/*M///////////////////////////////////////////////////////////////////////////////////////
//
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
//  By downloading, copying, installing or using the software you agree to this license.
//  If you do not agree to this license, do not download, install,
//  copy or use the software.
//
//
//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
//   * Redistribution's of source code must retain the above copyright notice,
//     this list of conditions and the following disclaimer.
//
//   * Redistribution's in binary form must reproduce the above copyright notice,
//     this list of conditions and the following disclaimer in the documentation
//     and/or other materials provided with the distribution.
//
//   * The name of the copyright holders may not be used to endorse or promote products
//     derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/

#ifndef OPENCV_CUDA_UTILITY_HPP
#define OPENCV_CUDA_UTILITY_HPP

#include "saturate_cast.hpp"
#include "datamov_utils.hpp"

/** @file
 * @deprecated Use @ref cudev instead.
 */

//! @cond IGNORED

namespace cv {
    namespace cuda {
        namespace device {
            struct CV_EXPORTS ThrustAllocator
                    {
                            typedef uchar value_type;
                            virtual ~ThrustAllocator();
                            virtual __device__ __host__ uchar* allocate(size_t numBytes) = 0;
                            virtual __device__ __host__ void deallocate(uchar* ptr, size_t numBytes) = 0;
                            static ThrustAllocator& getAllocator();
                            static void setAllocator(ThrustAllocator* allocator);
                    };
#define OPENCV_CUDA_LOG_WARP_SIZE        (5)
#define OPENCV_CUDA_WARP_SIZE            (1 << OPENCV_CUDA_LOG_WARP_SIZE)
#define OPENCV_CUDA_LOG_MEM_BANKS        ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla
#define OPENCV_CUDA_MEM_BANKS            (1 << OPENCV_CUDA_LOG_MEM_BANKS)

            ///////////////////////////////////////////////////////////////////////////////
            // swap

            template<typename T> void __device__
            __host__ __forceinline__
            swap(T
            & a,
            T &b
            ) {
            const T temp = a;
            a = b;
            b = temp;
        }

        ///////////////////////////////////////////////////////////////////////////////
        // Mask Reader

        struct SingleMask {
            explicit __host__ __device__

            __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}

            __host__ __device__

            __forceinline__ SingleMask(const SingleMask &mask_): mask(mask_.mask) {}

            __device__ __forceinline__

            bool operator()(int y, int x) const {
                return mask.ptr(y)[x] != 0;
            }

            PtrStepb mask;
        };

        struct SingleMaskChannels {
            __host__ __device__

            __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
                    : mask(mask_), channels(channels_) {}

            __host__ __device__

            __forceinline__ SingleMaskChannels(const SingleMaskChannels &mask_)
                    :mask(mask_.mask), channels(mask_.channels) {}

            __device__ __forceinline__

            bool operator()(int y, int x) const {
                return mask.ptr(y)[x / channels] != 0;
            }

            PtrStepb mask;
            int channels;
        };

        struct MaskCollection {
            explicit __host__ __device__

            __forceinline__ MaskCollection(PtrStepb *maskCollection_)
                    : maskCollection(maskCollection_) {}

            __device__ __forceinline__

            MaskCollection(const MaskCollection &masks_)
                    : maskCollection(masks_.maskCollection), curMask(masks_.curMask) {}

            __device__ __forceinline__

            void next() {
                curMask = *maskCollection++;
            }

            __device__ __forceinline__

            void setMask(int z) {
                curMask = maskCollection[z];
            }

            __device__ __forceinline__

            bool operator()(int y, int x) const {
                uchar val;
                return curMask.data == 0 ||
                       (ForceGlob<uchar>::Load(curMask.ptr(y), x, val), (val != 0));
            }

            const PtrStepb *maskCollection;
            PtrStepb curMask;
        };

        struct WithOutMask {
            __host__ __device__

            __forceinline__ WithOutMask() {}

            __host__ __device__

            __forceinline__ WithOutMask(const WithOutMask &) {}

            __device__ __forceinline__

            void next() const {
            }

            __device__ __forceinline__

            void setMask(int) const {
            }

            __device__ __forceinline__

            bool operator()(int, int) const {
                return true;
            }

            __device__ __forceinline__

            bool operator()(int, int, int) const {
                return true;
            }

            static __device__ __forceinline__

            bool check(int, int) {
                return true;
            }

            static __device__ __forceinline__

            bool check(int, int, int) {
                return true;
            }
        };

        ///////////////////////////////////////////////////////////////////////////////
        // Solve linear system

        // solve 2x2 linear system Ax=b
        template<typename T> __device__ __forceinline__

        bool solve2x2(const T A[2][2], const T b[2], T x[2]) {
            T det = A[0][0] * A[1][1] - A[1][0] * A[0][1];

            if (det != 0) {
                double invdet = 1.0 / det;

                x[0] = saturate_cast<T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1]));

                x[1] = saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0]));

                return true;
            }

            return false;
        }

        // solve 3x3 linear system Ax=b
        template<typename T> __device__ __forceinline__

        bool solve3x3(const T A[3][3], const T b[3], T x[3]) {
            T det = A[0][0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1])
                    - A[0][1] * (A[1][0] * A[2][2] - A[1][2] * A[2][0])
                    + A[0][2] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]);

            if (det != 0) {
                double invdet = 1.0 / det;

                x[0] = saturate_cast<T>(invdet *
                                        (b[0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) -
                                         A[0][1] * (b[1] * A[2][2] - A[1][2] * b[2]) +
                                         A[0][2] * (b[1] * A[2][1] - A[1][1] * b[2])));

                x[1] = saturate_cast<T>(invdet *
                                        (A[0][0] * (b[1] * A[2][2] - A[1][2] * b[2]) -
                                         b[0] * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) +
                                         A[0][2] * (A[1][0] * b[2] - b[1] * A[2][0])));

                x[2] = saturate_cast<T>(invdet *
                                        (A[0][0] * (A[1][1] * b[2] - b[1] * A[2][1]) -
                                         A[0][1] * (A[1][0] * b[2] - b[1] * A[2][0]) +
                                         b[0] * (A[1][0] * A[2][1] - A[1][1] * A[2][0])));

                return true;
            }

            return false;
        }
    }
}
} // namespace cv { namespace cuda { namespace cudev

//! @endcond

#endif // OPENCV_CUDA_UTILITY_HPP
